/*
 * Copyright 1993-2010 NVIDIA Corporation.  All rights reserved.
 *
 * NVIDIA Corporation and its licensors retain all intellectual property and 
 * proprietary rights in and to this software and related documentation. 
 * Any use, reproduction, disclosure, or distribution of this software 
 * and related documentation without an express license agreement from
 * NVIDIA Corporation is strictly prohibited.
 *
 * Please refer to the applicable NVIDIA end user license agreement (EULA) 
 * associated with this source code for terms and conditions that govern 
 * your use of this NVIDIA software.
 * 
 */



#ifndef QUASIRANDOMGENERATOR_KERNEL_CUH
#define QUASIRANDOMGENERATOR_KERNEL_CUH


#include <stdio.h>
#include <stdlib.h>
#include <cutil_inline.h>
#include "realtype.h"
#include "tsp_common.h"



//Fast integer multiplication
#define MUL(a, b) __umul24(a, b)



////////////////////////////////////////////////////////////////////////////////
// Niederreiter quasirandom number generation kernel
////////////////////////////////////////////////////////////////////////////////
static __constant__ unsigned int c_Table[QRNG_DIMENSIONS][QRNG_RESOLUTION];

__device__ int cudaRandom(unsigned int seed, unsigned int N){
		float resultado = 0;
    unsigned int *dimBase = &c_Table[threadIdx.y][0];
    unsigned int      tid = MUL(blockDim.x, blockIdx.x) + threadIdx.x;
    unsigned int  threadN = MUL(blockDim.x, gridDim.x);

    for(unsigned int pos = tid; pos < N; pos += threadN){
        unsigned int result = 0;
        unsigned int data = seed + pos;

        for(int bit = 0; bit < QRNG_RESOLUTION; bit++, data >>= 1)
            if(data & 1) result ^= dimBase[bit];

        resultado = (float)(result + 1) * INT_SCALE;
    }

		return (int)(resultado*1000000);
}

static __global__ void tspKernel(unsigned int N, int NCiudades, int *w){
    unsigned int tid = MUL(blockDim.x, blockIdx.x) + threadIdx.x;

		int aux = tid;
		for(int i=0; i<NCiudades; i++)
		{
				if(i<NCiudades) aux  = cudaRandom( (unsigned int)aux, N );
				w[i] = aux;
		}
}

static __global__ void quasirandomGeneratorKernel(
    float *d_Output,
    unsigned int seed,
    unsigned int N
){
    unsigned int *dimBase = &c_Table[threadIdx.y][0];
    unsigned int      tid = MUL(blockDim.x, blockIdx.x) + threadIdx.x;
    unsigned int  threadN = MUL(blockDim.x, gridDim.x);

    for(unsigned int pos = tid; pos < N; pos += threadN){
        unsigned int result = 0;
        unsigned int data = seed + pos;

        for(int bit = 0; bit < QRNG_RESOLUTION; bit++, data >>= 1)
            if(data & 1) result ^= dimBase[bit];

        d_Output[MUL(threadIdx.y, N) + pos] = (float)(result + 1) * INT_SCALE;
    }
}

//Table initialization routine
static void initTableGPU(unsigned int tableCPU[QRNG_DIMENSIONS][QRNG_RESOLUTION]){
    cutilSafeCall( cudaMemcpyToSymbol(
        c_Table,
        tableCPU,
        QRNG_DIMENSIONS * QRNG_RESOLUTION * sizeof(unsigned int)
    ) );
}

//Host-side interface
static void quasirandomGeneratorGPU(float *d_Output, unsigned int seed, unsigned int N){
    dim3 threads(128, QRNG_DIMENSIONS);
    quasirandomGeneratorKernel<<<128, threads>>>(d_Output, seed, N);
    cutilCheckMsg("quasirandomGeneratorKernel() execution failed.\n");
}

//Host-side interface
static void tspGPU(unsigned int N, int NCiudades, int optimo, int *distancias, int *w, int *wPrima){
	 	dim3 threads(128, QRNG_DIMENSIONS);
	 	printf("Llego a la función tspGPU, en tsp_kernel.cuh\n");
		tspKernel<<<128, threads>>>(N, NCiudades,w);
}

#endif
